#include <CL/cl.hpp>
#include <iostream>

#ifdef __cplusplus
extern "C" {
#endif

#include <stdio.h>
#include <math.h>
#include <string.h>
#include <stdlib.h>

#ifdef __cplusplus
}
#endif

#include "support.h"

#define PROGRAM_FILE "xxx.cl"
#define KERNEL_FUNC "innerProductKer_1"
#define VSIZE 153
#define GBSZ  1024
#define WGSZ  4
#define PASS  50

template <typename T> void
RunTest(cl_device_id id, cl_context ctx, cl_command_queue queue, const char* pragmaText);
// Generate OpenCL kernel code based on benchmark type struct
//void generateKernel (ostringstream &oss, struct _benchmark_type &test, const char* tName, const char* header);

int main (int argc, char *argv[]) {

	int platform = 0;
	int device = 0; 
/*
        cl::Device     id    = ListDevicesAndGetDevice(platform, device);
        std::vector<cl::Device> ctxDevices;
        ctxDevices.push_back( id );
        cl::Context ctx( ctxDevices );
        cl::CommandQueue queue( ctx, id, CL_QUEUE_PROFILING_ENABLE );
*/

	std::vector<cl::Platform> platformList;
	cl::Platform::get(&platformList);

	cl_context_properties cpros[] = {
		CL_CONTEXT_PLATFORM,
		(cl_context_properties)(platformList[0])(),0
	};
	cl::Context context(CL_DEVICE_TYPE_CPU,cpros);

	typedef std::vector<cl::Device> deviceContainer;
	deviceContainer devices = context.getInfo<CL_CONTEXT_DEVICES>();

	cl::CommandQueue queue(context,devices[0],CL_QUEUE_PROFILING_ENABLE);

	cl_device_id devid = devices[0]();
	cl_context ctx = context();
	cl_command_queue cq = queue(); 

	std::string pragma = std::string("#pragma OPENCL EXTENSION cl_amd_fp64: enable");
	RunTest<float>(devid, ctx, cq, pragma.c_str());

return 0;
}

template <typename T> void
RunTest(cl_device_id id, cl_context ctx, cl_command_queue queue, const char* pragmaText)
{
	FILE *program_handle;
   	char *program_buffer, *program_log;
	size_t program_size, log_size;

	cl_int err=0;

        // unsigned long vec_length = VSIZE;
        int vec_length = VSIZE;
        T * vec_1 =  new T[VSIZE]; 
        T * vec_2 =  new T[VSIZE];
        T * host_reduction_buffer = new T[GBSZ];

        for (int i=0; i< VSIZE;i++) {
            vec_1[i] = T(1);
//            std::cout<<vec_1[i]<<std::endl;
            vec_2[i] = T(2);
        }
        for (int i=0; i<GBSZ; i++) {
            host_reduction_buffer[i] = 100;
//            std::cout<<host_reduction_buffer[i]<<std::endl;
	}
	cl_mem cl_vec_1, cl_vec_2, cl_reduction_buffer;

	cl_vec_1 = clCreateBuffer(ctx, CL_MEM_READ_ONLY,
                                 sizeof(T)*VSIZE, NULL, &err);
        CL_CHECK_ERROR(err);
	cl_vec_2 = clCreateBuffer(ctx, CL_MEM_READ_ONLY,
                                 sizeof(T)*VSIZE, NULL, &err);
        CL_CHECK_ERROR(err);
	cl_reduction_buffer = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
                                 sizeof(T)*GBSZ, NULL, &err);
        CL_CHECK_ERROR(err);
// Issue a copy to force device allocation
        err = clEnqueueWriteBuffer(queue, cl_vec_1, true, 0,
                               vec_length*sizeof(T), vec_1,
                               0, NULL, NULL);
        CL_CHECK_ERROR(err);
        err = clEnqueueWriteBuffer(queue, cl_vec_2, true, 0,
                               vec_length*sizeof(T), vec_2,
                               0, NULL, NULL);
        CL_CHECK_ERROR(err);

	/* Read program file and place content into buffer */
   	program_handle = fopen(PROGRAM_FILE, "r");
   	if(program_handle == NULL) {
      	perror("Couldn't find the program file");
      	exit(1);
   	}
   	fseek(program_handle, 0, SEEK_END);
   	program_size = ftell(program_handle);
   	rewind(program_handle);
   	program_buffer = (char*)calloc(program_size+1, sizeof(char));
   	fread(program_buffer, sizeof(char), program_size, program_handle);
   	fclose(program_handle);	

// segment from cl.hpp example
	cl::Program::Sources source(1,
		std::make_pair(program_buffer,strlen(program_buffer)));

 	cl::Program program_ = cl::Program(ctx, source);
//	program_.build(devices);
 	cl_program prg = program_(); 
/*
	cl::Kernel kernel(program_, "hello", &err);
  
	cl::Event event;
	cl::CommandQueue queue(context, devices[0], 0, &err);
		queue.enqueueNDRangeKernel(
		kernel, 
		cl::NullRange, 
		cl::NDRange(4,4),
		cl::NullRange,
		NULL,
		&event); 

	event.wait();
*/
/* Build program */
	std::cout<< "build program start"<<std::endl;	
	err = clBuildProgram(prg, 0, NULL, "-w", NULL, NULL);
        CL_CHECK_ERROR(err);
	if(err < 0) {

      /* Find size of log and print to std output */
	clGetProgramBuildInfo(prg, id, CL_PROGRAM_BUILD_LOG,
		0, NULL, &log_size);
	program_log = (char*) calloc(log_size+1, sizeof(char));
	clGetProgramBuildInfo(prg, id, CL_PROGRAM_BUILD_LOG,
		log_size+1, program_log, NULL);
	printf("%s\n", program_log);
	free(program_log);
	exit(1);
	}		
	std::cout<< "build program over"<<std::endl;

/* Create a kernel */
	std::cout <<"create kernel start"<<std::endl;
	cl_kernel kernel;
	kernel = clCreateKernel(prg, KERNEL_FUNC, &err);
        CL_CHECK_ERROR(err);
	std::cout<< "create kernel over "<< err <<std::endl;


/* Create kernel argument */
	std::cout << "create kernel arg start"<<std::endl;
	err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &cl_vec_1);
	CL_CHECK_ERROR(err);
	err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &cl_vec_2);
	CL_CHECK_ERROR(err);
	err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &cl_reduction_buffer);
	CL_CHECK_ERROR(err);
	err = clSetKernelArg(kernel, 3, sizeof(vec_length), &vec_length);
	CL_CHECK_ERROR(err);
	err = clSetKernelArg(kernel, 4, WGSZ*sizeof(T), NULL);
	CL_CHECK_ERROR(err);
	std::cout << "create kernel arg over "<< err <<std::endl;

/* Enqueue kernel */
	std::cout << "enqueue kernel start"<<std::endl;
	size_t global_size = GBSZ;
	size_t local_size = WGSZ;
        for (int j = 0; j<PASS; j++) {
	     err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size,
		                          &local_size, 0, NULL, NULL);
	     CL_CHECK_ERROR(err);
             err = clEnqueueBarrier(queue);
	     CL_CHECK_ERROR(err);
             err = clFinish(queue);
	     CL_CHECK_ERROR(err);

	     err = clEnqueueReadBuffer(queue, cl_reduction_buffer, CL_TRUE, 0,
		                       GBSZ*sizeof(T), host_reduction_buffer, 0, NULL, NULL);
	     CL_CHECK_ERROR(err);

             std::cout << "global size is "<<global_size<<" , local size is "<<local_size<<std::endl;
             for (int k = 0; k< GBSZ ;k++)
                  std::cout<< "host_reduction_buffer " <<k<< " is "<<host_reduction_buffer[k]<<std::endl; 
        }
	std::cout << "enqueue kernel  part over"<<std::endl;

/* Print the result */
// clean Program Objects
       // err = clReleaseKernel(kernel);
       // CL_CHECK_ERROR(err);
//	err = clReleaseProgram(prg);
//	CL_CHECK_ERROR(err);
// clean Mem obj
	clReleaseMemObject(cl_vec_1);
	clReleaseMemObject(cl_vec_2);
	clReleaseMemObject(cl_reduction_buffer);


        delete[] vec_1;
        delete[] vec_2;
        delete[] host_reduction_buffer;
}
